#include "space_object_kernel.cuh"
#include "cuda_info.hh"

#include <stdio.h>

#ifdef CUDA_VARIANT_CPU
void space_object_cuda_collide(SpaceObject::SO_Data *objects, SpaceObjectAtlas::SOA_Data *atlas, unsigned int count){
	for(int i = 0; i < count - 1; i++){
		SpaceObject::SO_Data obj1 = objects[i];
		SpaceObjectAtlas::SOA_Data obj_a1 = atlas[obj1.atlas_id];
		for(int j = i + 1; j < count - 1; j++){
			SpaceObject::SO_Data obj2 = objects[j];
			SpaceObjectAtlas::SOA_Data obj_a2 = atlas[obj2.atlas_id];

			float distance = (obj1.x - obj2.x) * (obj1.x - obj2.x)
					+ (obj1.y - obj2.y) * (obj1.y - obj2.y)
					+ (obj1.z - obj2.z) * (obj1.z - obj2.z);
			if(distance < (obj_a1.sphereRadiusSqr + obj_a2.sphereRadiusSqr)){
				objects[i].collide = true;
				objects[index].collide = true;
				objects[i].collide_id = index;
				objects[index].collide_id = i;
			}
		}
	}
}

#endif

#ifdef CUDA_VARIANT_GPU1
// CUDA time: 40.000: 1227.4 ms, 1235.9 ms
__global__ void space_object_kernel_collide(SpaceObject::SO_Data *objects, SpaceObjectAtlas::SOA_Data *atlas, unsigned int count){
	int index = threadIdx.x + blockIdx.x * blockDim.x;

	if(index < count){
	SpaceObject::SO_Data obj1 = objects[index];
	SpaceObjectAtlas::SOA_Data obj_a1 = atlas[obj1.atlas_id];
	for(int i = index + 1; i < count; i++){
		SpaceObject::SO_Data obj2 = objects[i];
		SpaceObjectAtlas::SOA_Data obj_a2 = atlas[obj2.atlas_id];
		float distance = (obj1.x - obj2.x) * (obj1.x - obj2.x)
				+ (obj1.y - obj2.y) * (obj1.y - obj2.y)
				+ (obj1.z - obj2.z) * (obj1.z - obj2.z);
		if(distance < (obj_a1.sphereRadiusSqr + obj_a2.sphereRadiusSqr)){
			objects[i].collide = true;
			objects[index].collide = true;
			objects[i].collide_id = index;
			objects[index].collide_id = i;
		}
	}
	}
}

void space_object_cuda_collide(SpaceObject::SO_Data *objects, SpaceObjectAtlas::SOA_Data *atlas, unsigned int count){
	unsigned int blocks = 1;
	unsigned int threads = 1;
	unsigned int total_threads = count - 1;
	if(total_threads < cuda_info->threadsInBlock){
		threads = total_threads;
	} else {
		threads = cuda_info->threadsInBlock;
		blocks = (total_threads - 1) / threads + 1;
	}

	space_object_kernel_collide<<<blocks,threads>>>(objects, atlas, count);

}
#endif

#ifdef CUDA_VARIANT_GPU2
// CUDA time: 40.000: 2633.3 ms, 2611.3 ms
__global__ void space_object_kernel_collide(SpaceObject::SO_Data *objects, SpaceObjectAtlas::SOA_Data *atlas, unsigned int count){
	int index = threadIdx.x + blockIdx.x * blockDim.x;
	if(index < count / 2){

	for(int i = 1; i < count; i++){
		int idx1 = i > index ? index : count - index - 1;
		int idx2 = i > index ? i : count - 1 - index + i;
		SpaceObject::SO_Data obj1 = objects[idx1];
		SpaceObjectAtlas::SOA_Data obj_a1 = atlas[obj1.atlas_id];

		SpaceObject::SO_Data obj2 = objects[idx2];
		SpaceObjectAtlas::SOA_Data obj_a2 = atlas[obj2.atlas_id];
		float distance = (obj1.x - obj2.x) * (obj1.x - obj2.x)
				+ (obj1.y - obj2.y) * (obj1.y - obj2.y)
				+ (obj1.z - obj2.z) * (obj1.z - obj2.z);
		if(distance < (obj_a1.sphereRadiusSqr + obj_a2.sphereRadiusSqr)){
			objects[idx1].collide = true;
			objects[idx2].collide = true;
			objects[idx1].collide_id = index;
			objects[idx2].collide_id = i;
		}
	}
	}
}

void space_object_cuda_collide(SpaceObject::SO_Data *objects, SpaceObjectAtlas::SOA_Data *atlas, unsigned int count){
	unsigned int blocks = 1;
	unsigned int threads = 1;
	unsigned int total_threads = count / 2;
	if(total_threads < cuda_info->threadsInBlock){
		threads = total_threads;
	} else {
		threads = cuda_info->threadsInBlock;
		blocks = (total_threads - 1) / threads + 1;
	}

	space_object_kernel_collide<<<blocks,threads>>>(objects, atlas, count);

}
#endif

#ifdef CUDA_VARIANT_GPU3
__global__ void space_object_kernel_prepare(SpaceObject::SO_Data *objects, SpaceObjectAtlas::SOA_Data *atlas, unsigned int count){
	int index = threadIdx.x + blockIdx.x * blockDim.x;
	if(index < count){
		objects[index].sqrRad = atlas[objects[index].atlas_id].sphereRadiusSqr;
	}
}
// CUDA time: 40.000: 1227.4 ms, 1235.9 ms
__global__ void space_object_kernel_collide(SpaceObject::SO_Data *objects, SpaceObjectAtlas::SOA_Data *atlas, unsigned int count){
	int index = threadIdx.x + blockIdx.x * blockDim.x;

	if(index < count){
	SpaceObject::SO_Data obj1 = objects[index];
	for(int i = index + 1; i < count; i++){
		SpaceObject::SO_Data obj2 = objects[i];
		float distance = (obj1.x - obj2.x) * (obj1.x - obj2.x)
				+ (obj1.y - obj2.y) * (obj1.y - obj2.y)
				+ (obj1.z - obj2.z) * (obj1.z - obj2.z);
		if(distance < (obj1.sqrRad + obj2.sqrRad)){
			objects[i].collide = true;
			objects[index].collide = true;
			objects[i].collide_id = index;
			objects[index].collide_id = i;
		}
	}
	}
}

void space_object_cuda_collide(SpaceObject::SO_Data *objects, SpaceObjectAtlas::SOA_Data *atlas, unsigned int count){
	unsigned int blocks = 1;
	unsigned int threads = 1;
	unsigned int total_threads = count - 1;
	if(total_threads < cuda_info->threadsInBlock){
		threads = total_threads;
	} else {
		threads = cuda_info->threadsInBlock;
		blocks = (total_threads - 1) / threads + 1;
	}

	space_object_kernel_collide<<<blocks,threads>>>(objects, atlas, count);

}

void space_object_cuda_prepare(SpaceObject::SO_Data *objects, SpaceObjectAtlas::SOA_Data *atlas, unsigned int count){
	unsigned int blocks = 1;
	unsigned int threads = 1;
	unsigned int total_threads = count - 1;
	if(total_threads < cuda_info->threadsInBlock){
		threads = total_threads;
	} else {
		threads = cuda_info->threadsInBlock;
		blocks = (total_threads - 1) / threads + 1;
	}

	space_object_kernel_prepare<<<blocks,threads>>>(objects, atlas, count);

}
#endif
